#pragma OPENCL EXTENSION cl_arm_printf : enable
#pragma OPENCL EXTENSION cl_arm_integer_dot_product_int8 : enable

__kernel void matmul_int8_lbrbobntnp(__global char * out_buf, __global char4 *lhs_buf,
                                     __global char4 *rhs_buf, int4 hwkt) {
    int height = hwkt.x, width = hwkt.y, kslice = hwkt.z/4;
    int h = get_global_id(0);
    int w = get_global_id(1);

    // 203 ms
    int acc = 0;
    for(int i=0; i<kslice; i+=4) {
        acc += arm_dot(lhs_buf[h*kslice + i + 0], rhs_buf[w*kslice + i + 0]);
        acc += arm_dot(lhs_buf[h*kslice + i + 1], rhs_buf[w*kslice + i + 1]);
        acc += arm_dot(lhs_buf[h*kslice + i + 2], rhs_buf[w*kslice + i + 2]);
        acc += arm_dot(lhs_buf[h*kslice + i + 3], rhs_buf[w*kslice + i + 3]);
    }
    out_buf[h*width + w] = convert_char(acc % 53);

    // 708 ms
    // int acc = 0;
    // for(int i=0; i<kslice; i+=4) {
    //     acc += lhs_buf[h*kslice + i + 0].x * rhs_buf[w*kslice + i + 0].x;
    //     acc += lhs_buf[h*kslice + i + 0].y * rhs_buf[w*kslice + i + 0].y;
    //     acc += lhs_buf[h*kslice + i + 0].z * rhs_buf[w*kslice + i + 0].z;
    //     acc += lhs_buf[h*kslice + i + 0].w * rhs_buf[w*kslice + i + 0].w;

    //     acc += lhs_buf[h*kslice + i + 1].x * rhs_buf[w*kslice + i + 1].x;
    //     acc += lhs_buf[h*kslice + i + 1].y * rhs_buf[w*kslice + i + 1].y;
    //     acc += lhs_buf[h*kslice + i + 1].z * rhs_buf[w*kslice + i + 1].z;
    //     acc += lhs_buf[h*kslice + i + 1].w * rhs_buf[w*kslice + i + 1].w;

    //     acc += lhs_buf[h*kslice + i + 2].x * rhs_buf[w*kslice + i + 2].x;
    //     acc += lhs_buf[h*kslice + i + 2].y * rhs_buf[w*kslice + i + 2].y;
    //     acc += lhs_buf[h*kslice + i + 2].z * rhs_buf[w*kslice + i + 2].z;
    //     acc += lhs_buf[h*kslice + i + 2].w * rhs_buf[w*kslice + i + 2].w;

    //     acc += lhs_buf[h*kslice + i + 3].x * rhs_buf[w*kslice + i + 3].x;
    //     acc += lhs_buf[h*kslice + i + 3].y * rhs_buf[w*kslice + i + 3].y;
    //     acc += lhs_buf[h*kslice + i + 3].z * rhs_buf[w*kslice + i + 3].z;
    //     acc += lhs_buf[h*kslice + i + 3].w * rhs_buf[w*kslice + i + 3].w;
    // }
    // out_buf[h*width + w] = convert_char(acc % 53);

    // int4 acc = 0;
    // for(int i=0; i<kslice; i+=4) {
    //     acc = mad24(convert_int4(lhs_buf[h*kslice + i + 0]), convert_int4(rhs_buf[w*kslice + i + 0]), acc);
    //     acc = mad24(convert_int4(lhs_buf[h*kslice + i + 1]), convert_int4(rhs_buf[w*kslice + i + 1]), acc);
    //     acc = mad24(convert_int4(lhs_buf[h*kslice + i + 2]), convert_int4(rhs_buf[w*kslice + i + 2]), acc);
    //     acc = mad24(convert_int4(lhs_buf[h*kslice + i + 3]), convert_int4(rhs_buf[w*kslice + i + 3]), acc);
    // }
    // out_buf[h*width + w] = convert_char((acc.x+acc.y+acc.z+acc.w) % 53);

    // float acc = 0;
    // for(int i=0; i<kslice; i+=4) {
    //     acc += dot(convert_float4(lhs_buf[h*kslice + i + 0]), convert_float4(rhs_buf[w*kslice + i + 0]));
    //     acc += dot(convert_float4(lhs_buf[h*kslice + i + 1]), convert_float4(rhs_buf[w*kslice + i + 1]));
    //     acc += dot(convert_float4(lhs_buf[h*kslice + i + 2]), convert_float4(rhs_buf[w*kslice + i + 2]));
    //     acc += dot(convert_float4(lhs_buf[h*kslice + i + 3]), convert_float4(rhs_buf[w*kslice + i + 3]));
    // }
    // out_buf[h*width + w] = convert_char((int)(acc) % 53);
}